
#include "cuda_runtime.h"
#include "cuda_runtime_api.h"
#include "rgba_to_greyscale.h"

__global__ void rgba_to_greyscaleKernel(uchar4 *const rgbaImage, unsigned char *const greyImage, int numRows, int numCols)
{
    //获取每个线程的ID方法一
    // const int idx = blockIdx.x*blockDim.x + threadIdx.x;
    // const int idy = blockIdx.y*blockDim.y + threadIdx.y;
    // const int id = idx * numCols + idy;
    //方法二
    const int id = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;

    if (id < numRows * numCols)
    {
        const unsigned char R = rgbaImage[id].x;
        const unsigned char G = rgbaImage[id].y;
        const unsigned char B = rgbaImage[id].z;
        greyImage[id] = .299f * R + .587f * G + .114 * B;
    }
}


int rgba_to_greyscale(int height, int width, unsigned char* data, CallbackFun callback)
{
    uchar4 *h_rgbaImage, *d_rgbaImage;
    unsigned char *h_greyImage, *d_greyImage;

    //将imageRGBA的地址传出去目的是将Mat转换成数组
    h_rgbaImage= (uchar4 *)data;
    int numPixels = width * height;

    //在Device 上开辟内存
    cudaMalloc(&d_rgbaImage, sizeof(uchar4) * numPixels);
    cudaMalloc(&d_greyImage, sizeof(unsigned char) * numPixels);
    cudaHostAlloc(&h_greyImage, sizeof(unsigned char) * numPixels, cudaHostAllocMapped);
    //将d_greyImage全都初始化成0
    cudaMemset(d_greyImage, 0, numPixels * sizeof(unsigned char));

    //为d_rgbaImage赋值
    cudaMemcpy(d_rgbaImage, h_rgbaImage, sizeof(uchar4) * numPixels, cudaMemcpyHostToDevice);
    const int thread = 16;
    const int grid = (height * width + thread - 1) / (thread * thread);
    const dim3 blockSize(thread, thread);
    const dim3 gridSize(grid);
    rgba_to_greyscaleKernel<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, height, width);

    //等待线程全部结束
    cudaDeviceSynchronize();

    //将结果传回Host上
    cudaMemcpy(h_greyImage, d_greyImage, sizeof(unsigned char) * height * width, cudaMemcpyDeviceToHost);
    if (callback)
        callback(height, width, h_greyImage);

    // cleanup();
    cudaFree(d_rgbaImage);
    cudaFree(d_rgbaImage);
    cudaFreeHost(h_greyImage);

    return 0;
}
